{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Quene   队列是提交到SYCL运行的命令组，与设备是多对一的关系"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "q.submit([&](handler& h){\n",
    "    // command\n",
    "});"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Kernel  内核类封装了用于在实例化命令组时在设备上执行代码的方法和数据。内核对象不是由用户显式构造的，而是在调用内核调度函数（如parallel_for）时构造的"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "q.submit([&](handler& h){\n",
    "    h.paraller_for(range<1>[N],[=](id<1> i){\n",
    "        A[i] = B[i] + C[i];\n",
    "    });\n",
    "});"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "内核的功能可以通过id，range，item访问，item是内核的一个实例"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "q.submit([&](handler& h){\n",
    "    h.paraller_for(range<1>[N],[=](item<1> item){\n",
    "        auto i = item.get_id();\n",
    "        auto R = item.get_range();\n",
    "        // command\n",
    "    });\n",
    "});"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "队列运行的位置有三种方式"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "1. 面向任何设备 queue()\n",
    "2. 针对预配置的设备 如queue(cpu_selector{})\n",
    "3. 针对特定设备 class custom_selector : public device_selector {int operator()(…… // Any logic you want! …\n",
    "    queue(custom_selector{});"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "应用程序作用域和命令组作用域：\n",
    "1. 在主机上执行的代码\n",
    "2. C++的全部功能在应用程序和命令组范围内可用\n",
    "\n",
    "内核作用域：\n",
    "1. 在设备上执行的代码。\n",
    "2. 在内核范围内接受的C++功能"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "ND-range Kernel 一般内核不能在硬件级别上进行性能优化，  而ND-range将任务可以分配到更小的工作组中，工作组中的每一个项就被安排到硬件中执行。功能通过nd_range和nd_item访问。nd_item为核函数的实例"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "h.parallel_for(nd_range<1>(range<1>(1024),range<1>(64)), [=](nd_item<1> item){\n",
    "    auto idx = item.get_global_id();\n",
    "    auto local_id = itm.get_local_id();\n",
    "    // command\n",
    "});\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "缓冲区buffer和访问器accessor\n",
    "\n",
    "缓冲区将数据封装在跨设备和主机的 SYCL 应用程序中。访问器是访问缓冲区数据的机制。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "示例代码："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "void dpcpp_code(int* a, int* b, int* c, int N) {\n",
    "  //Step 1: create a device queue\n",
    "  //(developer can specify a device type via device selector or use default selector)\n",
    "  auto R = range<1>(N);\n",
    "  queue q;\n",
    "  //Step 2: create buffers (represent both host and device memory)\n",
    "  buffer buf_a(a, R);\n",
    "  buffer buf_b(b, R);\n",
    "  buffer buf_c(c, R);\n",
    "  //Step 3: submit a command for (asynchronous) execution\n",
    "  q.submit([&](handler &h){\n",
    "  //Step 4: create buffer accessors to access buffer data on the device\n",
    "  accessor A(buf_a,h,read_only);\n",
    "  accessor B(buf_b,h,read_only);\n",
    "  accessor C(buf_c,h,write_only);\n",
    "  \n",
    "  //Step 5: send a kernel (lambda) for execution\n",
    "  h.parallel_for(range<1>(N), [=](auto i){\n",
    "    //Step 6: write a kernel\n",
    "    //Kernel invocations are executed in parallel\n",
    "    //Kernel is invoked for each element of the range\n",
    "    //Kernel invocation has access to the invocation id\n",
    "    C[i] = A[i] + B[i];\n",
    "    });\n",
    "  });\n",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "若两个内核使用相同缓冲区则第二个内核需要等第一个内核使用完毕之后才能使用缓冲区"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <CL/sycl.hpp>\n",
    "\n",
    "constexpr int num=16;\n",
    "using namespace sycl;\n",
    "\n",
    "  int main() {\n",
    "  auto R = range<1>{ num };\n",
    "  //Create Buffers A and B\n",
    "  buffer<int> A{ R }, B{ R };\n",
    "  //Create a device queue\n",
    "  queue Q;\n",
    "  //Submit Kernel 1\n",
    "  Q.submit([&](handler& h) {\n",
    "    //Accessor for buffer A\n",
    "    accessor out(A,h,write_only);\n",
    "    h.parallel_for(R, [=](auto idx) {\n",
    "      out[idx] = idx[0]; }); });\n",
    "  //Submit Kernel 2\n",
    "  Q.submit([&](handler& h) {\n",
    "    //This task will wait till the first queue is complete\n",
    "    accessor out(A,h,write_only);\n",
    "    h.parallel_for(R, [=](auto idx) {\n",
    "      out[idx] += idx[0]; }); });\n",
    "  //Submit Kernel 3\n",
    "  Q.submit([&](handler& h) { \n",
    "    //Accessor for Buffer B\n",
    "    accessor out(B,h,write_only);\n",
    "    h.parallel_for(R, [=](auto idx) {\n",
    "      out[idx] = idx[0]; }); });\n",
    "  //Submit task 4\n",
    "  Q.submit([&](handler& h) {\n",
    "   //This task will wait till kernel 2 and 3 are complete\n",
    "   accessor in (A,h,read_only);\n",
    "   accessor inout(B,h);\n",
    "  h.parallel_for(R, [=](auto idx) {\n",
    "    inout[idx] *= in[idx]; }); }); \n",
    "      \n",
    " // And the following is back to device code\n",
    " host_accessor result(B,read_only);\n",
    "  for (int i=0; i<num; ++i)\n",
    "    std::cout << result[i] << \"\\n\";      \n",
    "  return 0;\n",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "主机访问器  \n",
    "\n",
    "在命令组的作用域之外创建的，并且它授予访问权限的数据将在主机上可用。它们用于通过构造主机访问器对象将数据同步回主机。缓冲区销毁是将数据同步回主机的另一种方法。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <CL/sycl.hpp>\n",
    "using namespace sycl;\n",
    "\n",
    "int main() {\n",
    "  constexpr int N = 16;\n",
    "  auto R = range<1>(N);\n",
    "  std::vector<int> v(N, 10);\n",
    "  queue q;\n",
    "  // Buffer takes ownership of the data stored in vector.  \n",
    "  buffer buf(v);\n",
    "  q.submit([&](handler& h) {\n",
    "    accessor a(buf,h);\n",
    "    h.parallel_for(R, [=](auto i) { a[i] -= 2; });\n",
    "  });\n",
    "  // Creating host accessor is a blocking call and will only return after all\n",
    "  // enqueued SYCL kernels that modify the same buffer in any queue completes\n",
    "  // execution and the data is available to the host via this host accessor.\n",
    "  host_accessor b(buf,read_only);\n",
    "  for (int i = 0; i < N; i++) std::cout << b[i] << \" \";\n",
    "  return 0;\n",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "缓冲区销毁示例\n",
    "\n",
    "缓冲区创建发生在单独的函数范围内。当执行超出此函数范围时，将调用缓冲区析构函数，该解析函数放弃数据的所有权并将数据复制回主机内存。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <CL/sycl.hpp>\n",
    "constexpr int N = 16;\n",
    "using namespace sycl;\n",
    "\n",
    "// Buffer creation happens within a separate function scope.\n",
    "void dpcpp_code(std::vector<int> &v, queue &q) {\n",
    "  auto R = range<1>(N);\n",
    "  buffer buf(v);\n",
    "  q.submit([&](handler &h) {\n",
    "    accessor a(buf,h);\n",
    "    h.parallel_for(R, [=](auto i) { a[i] -= 2; });\n",
    "  });\n",
    "}\n",
    "int main() {\n",
    "  std::vector<int> v(N, 10);\n",
    "  queue q;\n",
    "  dpcpp_code(v, q);\n",
    "  // When execution advances beyond this function scope, buffer destructor is\n",
    "  // invoked which relinquishes the ownership of data and copies back the data to\n",
    "  // the host memory.\n",
    "  for (int i = 0; i < N; i++) std::cout << v[i] << \" \";\n",
    "  return 0;\n",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "自定义设备选择器"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <iostream>\n",
    "using namespace sycl;\n",
    "class my_device_selector : public device_selector {\n",
    "public:\n",
    "    my_device_selector(std::string vendorName) : vendorName_(vendorName){};\n",
    "    int operator()(const device& dev) const override {\n",
    "        int rating = 0;\n",
    "        //We are querying for the custom device specific to a Vendor and if it is a GPU device we\n",
    "        //are giving the highest rating as 3 . The second preference is given to any GPU device and the third preference is given to\n",
    "        //CPU device.\n",
    "        if (dev.is_gpu() & (dev.get_info<info::device::name>().find(vendorName_) != std::string::npos))\n",
    "            rating = 3;\n",
    "        else if (dev.is_gpu()) rating = 2;\n",
    "        else if (dev.is_cpu()) rating = 1;\n",
    "        return rating;\n",
    "    };\n",
    "    \n",
    "private:\n",
    "    std::string vendorName_;\n",
    "};\n",
    "int main() {\n",
    "    //pass in the name of the vendor for which the device you want to query \n",
    "    std::string vendor_name = \"Intel\";\n",
    "    //std::string vendor_name = \"AMD\";\n",
    "    //std::string vendor_name = \"Nvidia\";\n",
    "    my_device_selector selector(vendor_name);\n",
    "    queue q(selector);\n",
    "    std::cout << \"Device: \"\n",
    "    << q.get_device().get_info<info::device::name>() << \"\\n\";\n",
    "    return 0;\n",
    "}\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "复数乘法"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "// complex.hpp\n",
    "#include <iostream>\n",
    "#include <vector>\n",
    "using namespace std;\n",
    "class Complex2 {\n",
    " private:\n",
    "  int m_real_, m_imag_;\n",
    "\n",
    " public:\n",
    "  Complex2() {\n",
    "    m_real_ = 0;\n",
    "    m_imag_ = 0;\n",
    "  }\n",
    "  Complex2(int x, int y) {\n",
    "    m_real_ = x;\n",
    "    m_imag_ = y;\n",
    "  }\n",
    "\n",
    "  // Overloading the  != operator\n",
    "  friend bool operator!=(const Complex2& a, const Complex2& b) {\n",
    "    return (a.m_real_ != b.m_real_) || (a.m_imag_ != b.m_imag_);\n",
    "  }\n",
    "\n",
    "  // The function performs Complex number multiplication and returns a Complex2\n",
    "  // object.\n",
    "  Complex2 complex_mul(const Complex2& obj) const {\n",
    "    return Complex2(((m_real_ * obj.m_real_) - (m_imag_ * obj.m_imag_)),\n",
    "                    ((m_real_ * obj.m_imag_) + (m_imag_ * obj.m_real_)));\n",
    "  }\n",
    "\n",
    "  // Overloading the ostream operator to print the objects of the Complex2\n",
    "  // object\n",
    "  friend ostream& operator<<(ostream& out, const Complex2& obj) {\n",
    "    out << \"(\" << obj.m_real_ << \" : \" << obj.m_imag_ << \"i)\";\n",
    "    return out;\n",
    "  }\n",
    "};"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <CL/sycl.hpp>\n",
    "#include <iomanip>\n",
    "#include <vector>\n",
    "// dpc_common.hpp can be found in the dev-utilities include folder.\n",
    "// e.g., $ONEAPI_ROOT/dev-utilities/<version>/include/dpc_common.hpp\n",
    "#include \"dpc_common.hpp\"\n",
    "#include \"Complex.hpp\"\n",
    "\n",
    "using namespace sycl;\n",
    "using namespace std;\n",
    "\n",
    "// Number of complex numbers passing to the SYCL code\n",
    "static const int num_elements = 10000;\n",
    "\n",
    "class CustomDeviceSelector : public device_selector {\n",
    " public:\n",
    "  CustomDeviceSelector(std::string vendorName) : vendorName_(vendorName){};\n",
    "  int operator()(const device &dev) const override {\n",
    "    int device_rating = 0;\n",
    "    //We are querying for the custom device specific to a Vendor and if it is a GPU device we\n",
    "    //are giving the highest rating as 3 . The second preference is given to any GPU device and the third preference is given to\n",
    "    //CPU device. \n",
    "    if (dev.is_gpu() & (dev.get_info<info::device::name>().find(vendorName_) !=\n",
    "                        std::string::npos))\n",
    "      device_rating = 3;\n",
    "    else if (dev.is_gpu())\n",
    "      device_rating = 2;\n",
    "    else if (dev.is_cpu())\n",
    "      device_rating = 1;\n",
    "    return device_rating;\n",
    "  };\n",
    "\n",
    " private:\n",
    "  std::string vendorName_;\n",
    "};\n",
    "\n",
    "// in_vect1 and in_vect2 are the vectors with num_elements complex nubers and\n",
    "// are inputs to the parallel function\n",
    "void DpcppParallel(queue &q, std::vector<Complex2> &in_vect1,\n",
    "                   std::vector<Complex2> &in_vect2,\n",
    "                   std::vector<Complex2> &out_vect) {\n",
    "  auto R = range(in_vect1.size());\n",
    "  if (in_vect2.size() != in_vect1.size() || out_vect.size() != in_vect1.size()){ \n",
    "    std::cout << \"ERROR: Vector sizes do not  match\"<< \"\\n\";\n",
    "    return;\n",
    "  }\n",
    "  // Setup input buffers\n",
    "  buffer bufin_vect1(in_vect1);\n",
    "  buffer bufin_vect2(in_vect2);\n",
    "\n",
    "  // Setup Output buffers \n",
    "  buffer bufout_vect(out_vect);\n",
    "\n",
    "  std::cout << \"Target Device: \"\n",
    "            << q.get_device().get_info<info::device::name>() << \"\\n\";\n",
    "  // Submit Command group function object to the queue\n",
    "  q.submit([&](auto &h) {\n",
    "    // Accessors set as read mode\n",
    "    accessor V1(bufin_vect1,h,read_only);\n",
    "    accessor V2(bufin_vect2,h,read_only);\n",
    "    // Accessor set to Write mode\n",
    "    accessor V3 (bufout_vect,h,write_only);\n",
    "    h.parallel_for(R, [=](auto i) {\n",
    "      V3[i] = V1[i].complex_mul(V2[i]);\n",
    "    });\n",
    "  });\n",
    "  q.wait_and_throw();\n",
    "}\n",
    "void DpcppScalar(std::vector<Complex2> &in_vect1,\n",
    "                 std::vector<Complex2> &in_vect2,\n",
    "                 std::vector<Complex2> &out_vect) {\n",
    "  if ((in_vect2.size() != in_vect1.size()) || (out_vect.size() != in_vect1.size())){\n",
    "    std::cout<<\"ERROR: Vector sizes do not match\"<<\"\\n\";\n",
    "    return;\n",
    "    }\n",
    "  for (int i = 0; i < in_vect1.size(); i++) {\n",
    "    out_vect[i] = in_vect1[i].complex_mul(in_vect2[i]);\n",
    "  }\n",
    "}\n",
    "// Compare the results of the two output vectors from parallel and scalar. They\n",
    "// should be equal\n",
    "int Compare(std::vector<Complex2> &v1, std::vector<Complex2> &v2) {\n",
    "  int ret_code = 1;\n",
    "  if(v1.size() != v2.size()){\n",
    "    ret_code = -1;\n",
    "  }\n",
    "  for (int i = 0; i < v1.size(); i++) {\n",
    "    if (v1[i] != v2[i]) {\n",
    "      ret_code = -1;\n",
    "      break;\n",
    "    }\n",
    "  }\n",
    "  return ret_code;\n",
    "}\n",
    "int main() {\n",
    "  // Declare your Input and Output vectors of the Complex2 class\n",
    "  vector<Complex2> input_vect1;\n",
    "  vector<Complex2> input_vect2;\n",
    "  vector<Complex2> out_vect_parallel;\n",
    "  vector<Complex2> out_vect_scalar;\n",
    "\n",
    "  for (int i = 0; i < num_elements; i++) {\n",
    "    input_vect1.push_back(Complex2(i + 2, i + 4));\n",
    "    input_vect2.push_back(Complex2(i + 4, i + 6));\n",
    "    out_vect_parallel.push_back(Complex2(0, 0));\n",
    "    out_vect_scalar.push_back(Complex2(0, 0));\n",
    "  }\n",
    "\n",
    "  // Initialize your Input and Output Vectors. Inputs are initialized as below.\n",
    "  // Outputs are initialized with 0\n",
    "  try {\n",
    "    // Pass in the name of the vendor for which the device you want to query\n",
    "    std::string vendor_name = \"Intel\";\n",
    "    // std::string vendor_name = \"AMD\";\n",
    "    // std::string vendor_name = \"Nvidia\";\n",
    "    // queue constructor passed exception handler\n",
    "    CustomDeviceSelector selector(vendor_name);\n",
    "    queue q(selector, dpc_common::exception_handler);\n",
    "    // Call the DpcppParallel with the required inputs and outputs\n",
    "    DpcppParallel(q, input_vect1, input_vect2, out_vect_parallel);\n",
    "  } catch (...) {\n",
    "    // some other exception detected\n",
    "    std::cout << \"Failure\" << \"\\n\";\n",
    "    std::terminate();\n",
    "  }\n",
    "\n",
    "  std::cout\n",
    "      << \"****************************************Multiplying Complex numbers \"\n",
    "         \"in Parallel********************************************************\"\n",
    "      << \"\\n\";\n",
    "  // Print the outputs of the Parallel function\n",
    "  int indices[]{0, 1, 2, 3, 4, (num_elements - 1)};\n",
    "  constexpr size_t indices_size = sizeof(indices) / sizeof(int);\n",
    "\n",
    "  for (int i = 0; i < indices_size; i++) {\n",
    "    int j = indices[i];\n",
    "    if (i == indices_size - 1) std::cout << \"...\\n\";\n",
    "    std::cout << \"[\" << j << \"] \" << input_vect1[j] << \" * \" << input_vect2[j]\n",
    "              << \" = \" << out_vect_parallel[j] << \"\\n\";\n",
    "  }\n",
    "  // Call the DpcppScalar function with the required input and outputs\n",
    "  DpcppScalar(input_vect1, input_vect2, out_vect_scalar);\n",
    "\n",
    "  // Compare the outputs from the parallel and the scalar functions. They should\n",
    "  // be equal\n",
    "\n",
    "  int ret_code = Compare(out_vect_parallel, out_vect_scalar);\n",
    "  if (ret_code == 1) {\n",
    "    std::cout << \"Complex multiplication successfully run on the device\"\n",
    "              << \"\\n\";\n",
    "  } else\n",
    "    std::cout\n",
    "        << \"*********************************************Verification Failed. Results are \"\n",
    "           \"not matched**************************\"\n",
    "        << \"\\n\";\n",
    "\n",
    "  return 0;\n",
    "}\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "使用 SYCL 缓冲区和访问器概念完成下面的编码练习：\n",
    "\n",
    "代码在主机上初始化了三个向量vector1\n",
    "\n",
    "核函数将向量1递增\n",
    "\n",
    "创建一个新的向量2并初始化值为20，并为它分配缓冲区\n",
    "\n",
    "为第二个缓冲区增加第二个访问器\n",
    "\n",
    "将递增修改为相加，计算两个向量之和\n",
    "\n",
    "内核代码将 递增 1。vector1\n",
    "\n",
    "创建一个新的秒并初始化为值 20。vector2\n",
    "\n",
    "为上述第二个向量创建 sycl 缓冲区\n",
    "\n",
    "在内核代码中，为第二个矢量缓冲区创建第二个访问器\n",
    "\n",
    "修改矢量增量以矢量添加，方法是将vector2vector1"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <CL/sycl.hpp>\n",
    "\n",
    "using namespace sycl;\n",
    "\n",
    "int main() {\n",
    "    const int N = 256;\n",
    "    \n",
    "    //# Initialize a vector and print values\n",
    "    std::vector<int> vector1(N, 10);\n",
    "    std::cout<<\"\\nInput Vector1: \";    \n",
    "    for (int i = 0; i < N; i++) std::cout << vector1[i] << \" \";\n",
    "\n",
    "    //# STEP 1 : Create second vector, initialize to 20 and print values\n",
    "    std::vector<int> vector2(N, 20);\n",
    "    std::cout<<\"\\nInput Vector2: \";\n",
    "    for (int i = 0; i < N; i++) std::cout << vector2[i] << \" \"; \n",
    "    \n",
    "    //# Create Buffer\n",
    "    buffer vector1_buffer(vector1);\n",
    "    \n",
    "    //# STEP 2 : Create buffer for second vector \n",
    "    buffer vector2_buffer(vector2);\n",
    "\n",
    "    //# Submit task to add vector\n",
    "    queue q;\n",
    "    q.submit([&](handler &h) {\n",
    "      //# Create accessor for vector1_buffer\n",
    "      accessor vector1_accessor (vector1_buffer,h);\n",
    "      \n",
    "      //# STEP 3 - add second accessor for second buffer\n",
    "      accessor vector2_accessor (vector2_buffer,h);\n",
    "      \n",
    "      h.parallel_for(range<1>(N), [=](id<1> index) {\n",
    "\n",
    "        //# STEP 4 : Modify the code below to add the second vector to first one\n",
    "        vector1_accessor[index] += vector2_accessor[index];\n",
    "      });\n",
    "   });\n",
    "\n",
    " \n",
    "  //# Create a host accessor to copy data from device to host\n",
    "  host_accessor h_a(vector1_buffer,read_only);\n",
    "\n",
    "  //# Print Output values \n",
    "  std::cout<<\"\\nOutput Values: \";\n",
    "  for (int i = 0; i < N; i++) std::cout<< vector1[i] << \" \";\n",
    "  std::cout<<\"\\n\";\n",
    "\n",
    "  return 0;\n",
    "}\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": []
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": []
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3.9.12 ('base')",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.9.12"
  },
  "orig_nbformat": 4,
  "vscode": {
   "interpreter": {
    "hash": "1fb013cd28f8030a2eaa4fd1a9aeaafac50cba64f0c2f90a55e4f5253fcc3bd6"
   }
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
